





<!DOCTYPE html>
<html class="writer-html5" lang="zh-CN" >
<head>
  <meta charset="utf-8">
  
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <title>Auto-scheduling a Convolution Layer for GPU &mdash; tvm 0.8.dev1982 文档</title>
  

  
  <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/css/bootstrap.min.css" integrity="sha384-Gn5384xqQ1aoWXA+058RXPxPg6fy4IWvTNh0E263XmFcJlSAwiGgFAW/dAiS6JXm" crossorigin="anonymous">
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/gallery.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/tlcpack_theme.css" type="text/css" />

  
  
    <link rel="shortcut icon" href="../../_static/tvm-logo-square.png"/>
  

  
  
  
  
    
      <script type="text/javascript" id="documentation_options" data-url_root="../../" src="../../_static/documentation_options.js"></script>
        <script data-url_root="../../" id="documentation_options" src="../../_static/documentation_options.js"></script>
        <script src="../../_static/jquery.js"></script>
        <script src="../../_static/underscore.js"></script>
        <script src="../../_static/doctools.js"></script>
        <script src="../../_static/translations.js"></script>
    
    <script type="text/javascript" src="../../_static/js/theme.js"></script>

    
    <script type="text/javascript" src="../../_static/js/tlcpack_theme.js"></script>
    <link rel="index" title="索引" href="../../genindex.html" />
    <link rel="search" title="搜索" href="../../search.html" />
    <link rel="next" title="Auto-scheduling a Neural Network for x86 CPU" href="tune_network_x86.html" />
    <link rel="prev" title="Use AutoScheduler for Template-Free Scheduling" href="index.html" /> 
</head>

<body class="wy-body-for-nav">

   
  <div class="wy-grid-for-nav">
    
    
<header class="header">
    <div class="innercontainer">
      <div class="headerInner d-flex justify-content-between align-items-center">
          <div class="headerLogo">
               <a href="https://tvm.apache.org/"><img src=https://tvm.apache.org/assets/images/logo.svg alt="logo"></a>
          </div>

          <div id="headMenu" class="headerNav">
            <button type="button" id="closeHeadMenu" class="navCloseBtn"><img src="../../_static/img/close-icon.svg" alt="Close"></button>
             <ul class="nav">
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/community>Community</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/download>Download</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/vta>VTA</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/blog>Blog</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/docs>Docs</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmconf.org>Conference</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://github.com/apache/tvm/>Github</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmchinese.github.io/declaration_zh_CN.html>About-Translators</a>
                </li>
             </ul>
               <div class="responsivetlcdropdown">
                 <button type="button" class="btn-link">
                   ASF
                 </button>
                 <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                 </ul>
               </div>
          </div>
            <div class="responsiveMenuIcon">
              <button type="button" id="menuBtn" class="btn-menu"><img src="../../_static/img/menu-icon.svg" alt="Menu Icon"></button>
            </div>

            <div class="tlcDropdown">
              <div class="dropdown">
                <button type="button" class="btn-link dropdown-toggle" data-toggle="dropdown" aria-haspopup="true" aria-expanded="false">
                  ASF
                </button>
                <div class="dropdown-menu dropdown-menu-right">
                  <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                  </ul>
                </div>
              </div>
          </div>
       </div>
    </div>
 </header>
 
    <nav data-toggle="wy-nav-shift" class="wy-nav-side fixed">
      <div class="wy-side-scroll">
        <div class="wy-side-nav-search" >
          

          
            <a href="../../index.html">
          

          
            
            <img src="../../_static/tvm-logo-small.png" class="logo" alt="Logo"/>
          
          </a>

          
            
            
                <div class="version">
                  0.8.dev1982
                </div>
            
          

          
<div role="search">
  <form id="rtd-search-form" class="wy-form" action="../../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
    <input type="hidden" name="check_keywords" value="yes" />
    <input type="hidden" name="area" value="default" />
  </form>
</div>

          
        </div>

        
        <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
          
            
            
              
            
            
              <p class="caption" role="heading"><span class="caption-text">如何开始</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../install/index.html">安装 TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../contribute/index.html">贡献者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">用户引导</span></p>
<ul class="current">
<li class="toctree-l1"><a class="reference internal" href="../../tutorial/index.html">User Tutorial</a></li>
<li class="toctree-l1 current"><a class="reference internal" href="../index.html">How To Guides</a><ul class="current">
<li class="toctree-l2"><a class="reference internal" href="../compile_models/index.html">编译深度学习模型</a></li>
<li class="toctree-l2"><a class="reference internal" href="../deploy/index.html">TVM 部署模型和集成</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_relay/index.html">Work With Relay</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_schedules/index.html">Work With Tensor Expression and Schedules</a></li>
<li class="toctree-l2"><a class="reference internal" href="../optimize_operators/index.html">优化张量算子</a></li>
<li class="toctree-l2"><a class="reference internal" href="../tune_with_autotvm/index.html">Auto-Tune with Templates and AutoTVM</a></li>
<li class="toctree-l2 current"><a class="reference internal" href="index.html">Use AutoScheduler for Template-Free Scheduling</a><ul class="current">
<li class="toctree-l3 current"><a class="current reference internal" href="#">Auto-scheduling a Convolution Layer for GPU</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#define-the-computation">Define the computation</a></li>
<li class="toctree-l4"><a class="reference internal" href="#create-the-search-task">Create the search task</a></li>
<li class="toctree-l4"><a class="reference internal" href="#run-the-search">Run the search</a></li>
<li class="toctree-l4"><a class="reference internal" href="#check-correctness-and-evaluate-performance">Check correctness and evaluate performance</a></li>
<li class="toctree-l4"><a class="reference internal" href="#using-the-record-file">Using the record file</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="tune_network_x86.html">Auto-scheduling a Neural Network for x86 CPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_network_cuda.html">Auto-scheduling a Neural Network for NVIDIA GPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_network_arm.html">Auto-scheduling a Neural Network for ARM CPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_network_mali.html">Auto-scheduling a Neural Network for mali GPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_sparse_x86.html">Auto-scheduling Sparse Matrix Multiplication on CPU with Custom Sketch Rule</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_microtvm/index.html">Work With microTVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../extend_tvm/index.html">Extend TVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../profile/index.html">Profile Models</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../errors.html">Handle TVM Errors</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../faq.html">常见提问</a></li>
</ul>
</li>
</ul>
<p class="caption" role="heading"><span class="caption-text">开发者引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../dev/tutorial/index.html">Developer Tutorial</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../dev/how_to/how_to.html">开发者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">架构指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../arch/index.html">Design and Architecture</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">主题引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../topic/microtvm/index.html">microTVM：裸机使用TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../topic/vta/index.html">VTA: Versatile Tensor Accelerator</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">参考指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../reference/langref/index.html">语言参考</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/python/index.html">Python API</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/links.html">Other APIs</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/publications.html">Publications</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../genindex.html">索引</a></li>
</ul>

            
          
        </div>
        
      </div>
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
      
      <nav class="wy-nav-top" aria-label="top navigation" data-toggle="wy-nav-top">
        
            <div class="togglemenu">

            </div>
            <div class="nav-content">
              <!-- tvm -->
              Table of content
            </div>
        
      </nav>


      <div class="wy-nav-content">
        
        <div class="rst-content">
        

          




















<div role="navigation" aria-label="breadcrumbs navigation">

  <ul class="wy-breadcrumbs">
    
      <li><a href="../../index.html">Docs</a> <span class="br-arrow">></span></li>
        
          <li><a href="../index.html">How To Guides</a> <span class="br-arrow">></span></li>
        
          <li><a href="index.html">Use AutoScheduler for Template-Free Scheduling</a> <span class="br-arrow">></span></li>
        
      <li>Auto-scheduling a Convolution Layer for GPU</li>
    
    
      <li class="wy-breadcrumbs-aside">
        
            
            <a href="../../_sources/how_to/tune_with_autoscheduler/tune_conv2d_layer_cuda.rst.txt" rel="nofollow"> <img src="../../_static//img/source.svg" alt="viewsource"/></a>
          
        
      </li>
    
  </ul>

  
  <hr/>
</div>
          <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
           <div itemprop="articleBody">
            
  <div class="sphx-glr-download-link-note admonition note">
<p class="admonition-title">注解</p>
<p>Click <a class="reference internal" href="#sphx-glr-download-how-to-tune-with-autoscheduler-tune-conv2d-layer-cuda-py"><span class="std std-ref">here</span></a> to download the full example code</p>
</div>
<div class="sphx-glr-example-title section" id="auto-scheduling-a-convolution-layer-for-gpu">
<span id="auto-scheduler-conv-gpu"></span><span id="sphx-glr-how-to-tune-with-autoscheduler-tune-conv2d-layer-cuda-py"></span><h1>Auto-scheduling a Convolution Layer for GPU<a class="headerlink" href="#auto-scheduling-a-convolution-layer-for-gpu" title="永久链接至标题">¶</a></h1>
<p><strong>作者</strong>: <a class="reference external" href="https://github.com/merrymercy">Lianmin Zheng</a>,             <a class="reference external" href="https://github.com/jcf94/">Chengfan Jia</a></p>
<p>This is a tutorial on how to use the auto-scheduler for GPUs.</p>
<p>Different from the template-based <a class="reference internal" href="../tune_with_autotvm/index.html#tutorials-autotvm-sec"><span class="std std-ref">autotvm</span></a> which relies on
manual templates to define the search space, the auto-scheduler does not require any templates.
Users only need to write the computation declaration without any schedule commands or templates.
The auto-scheduler can automatically generate a large search space and
find a good schedule in the space.</p>
<p>We use a convolution layer as an example in this tutorial.</p>
<p>Note that this tutorial will not run on Windows or recent versions of macOS. To
get it to run, you will need to wrap the body of this tutorial in a <code class="code docutils literal notranslate"><span class="pre">if</span>
<span class="pre">__name__</span> <span class="pre">==</span> <span class="pre">&quot;__main__&quot;:</span></code> block.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">os</span>

<span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span>
<span class="kn">import</span> <span class="nn">tvm</span>
<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">te</span><span class="p">,</span> <span class="n">auto_scheduler</span><span class="p">,</span> <span class="n">topi</span>
<span class="kn">from</span> <span class="nn">tvm.topi.testing</span> <span class="k">import</span> <span class="n">conv2d_nchw_python</span>
</pre></div>
</div>
<div class="section" id="define-the-computation">
<h2>Define the computation<a class="headerlink" href="#define-the-computation" title="永久链接至标题">¶</a></h2>
<p>To begin with, let us define the computation of a convolution layer.
The function should return the list of input/output tensors.
From these tensors, the auto-scheduler can get the whole computational graph.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nd">@auto_scheduler</span><span class="o">.</span><span class="n">register_workload</span>
<span class="k">def</span> <span class="nf">conv2d_layer</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">stride</span><span class="p">,</span> <span class="n">padding</span><span class="p">):</span>
    <span class="n">data</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">N</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;data&quot;</span><span class="p">)</span>
    <span class="n">kernel</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;kernel&quot;</span><span class="p">)</span>
    <span class="n">bias</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="mi">1</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;bias&quot;</span><span class="p">)</span>
    <span class="n">conv</span> <span class="o">=</span> <span class="n">topi</span><span class="o">.</span><span class="n">nn</span><span class="o">.</span><span class="n">conv2d_nchw</span><span class="p">(</span><span class="n">data</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">stride</span><span class="p">,</span> <span class="n">padding</span><span class="p">,</span> <span class="n">dilation</span><span class="o">=</span><span class="mi">1</span><span class="p">,</span> <span class="n">out_dtype</span><span class="o">=</span><span class="s2">&quot;float32&quot;</span><span class="p">)</span>
    <span class="n">out</span> <span class="o">=</span> <span class="n">topi</span><span class="o">.</span><span class="n">nn</span><span class="o">.</span><span class="n">relu</span><span class="p">(</span><span class="n">conv</span> <span class="o">+</span> <span class="n">bias</span><span class="p">)</span>
    <span class="k">return</span> <span class="p">[</span><span class="n">data</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">bias</span><span class="p">,</span> <span class="n">out</span><span class="p">]</span>
</pre></div>
</div>
</div>
<div class="section" id="create-the-search-task">
<h2>Create the search task<a class="headerlink" href="#create-the-search-task" title="永久链接至标题">¶</a></h2>
<p>We then create a search task for the last convolution layer in the resnet.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">target</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span><span class="p">(</span><span class="s2">&quot;cuda&quot;</span><span class="p">)</span>

<span class="c1"># Use the last layer in ResNet-50</span>
<span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span> <span class="o">=</span> <span class="mi">1</span><span class="p">,</span> <span class="mi">7</span><span class="p">,</span> <span class="mi">7</span><span class="p">,</span> <span class="mi">512</span><span class="p">,</span> <span class="mi">512</span><span class="p">,</span> <span class="mi">3</span><span class="p">,</span> <span class="mi">3</span><span class="p">,</span> <span class="p">(</span><span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">),</span> <span class="p">(</span><span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">)</span>
<span class="n">task</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">SearchTask</span><span class="p">(</span>
    <span class="n">func</span><span class="o">=</span><span class="n">conv2d_layer</span><span class="p">,</span> <span class="n">args</span><span class="o">=</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span><span class="p">),</span> <span class="n">target</span><span class="o">=</span><span class="n">target</span>
<span class="p">)</span>

<span class="c1"># Inspect the computational graph</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Computational DAG:&quot;</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">task</span><span class="o">.</span><span class="n">compute_dag</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Computational DAG:
data = PLACEHOLDER [1, 512, 7, 7]
pad_temp(i0, i1, i2, i3) = tir.if_then_else(((((i2 &gt;= 1) &amp;&amp; (i2 &lt; 8)) &amp;&amp; (i3 &gt;= 1)) &amp;&amp; (i3 &lt; 8)), data[i0, i1, (i2 - 1), (i3 - 1)], 0f)
kernel = PLACEHOLDER [512, 512, 3, 3]
compute(nn, ff, yy, xx) += (pad_temp[nn, rc, (yy + ry), (xx + rx)]*kernel[ff, rc, ry, rx])
bias = PLACEHOLDER [1, 512, 1, 1]
T_add(ax0, ax1, ax2, ax3) = (compute[ax0, ax1, ax2, ax3] + bias[ax0, ax1, 0, 0])
compute(i0, i1, i2, i3) = max(T_add[i0, i1, i2, i3], 0f)
</pre></div>
</div>
<p>Next, we set parameters for the auto-scheduler. These parameters
mainly specify how we do the measurement during the search.</p>
<ul class="simple">
<li><p><code class="code docutils literal notranslate"><span class="pre">measure_ctx</span></code> launches a different process for measurement to
provide isolation. It can protect the master process from GPU crashes
during measurement and avoid other runtime conflicts.</p></li>
<li><p><code class="code docutils literal notranslate"><span class="pre">min_repeat_ms</span></code> defines the minimum duration of one “repeat” in every measurement.
This can warmup the GPU, which is necessary to get accurate measurement results.
Typically, we recommend a value &gt;= 300 ms.</p></li>
<li><p><code class="code docutils literal notranslate"><span class="pre">num_measure_trials</span></code> is the number of measurement trials we can use during the search.
We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a
good value for the search to converge. You can do more trials according to your time budget.</p></li>
<li><p>In addition, we use <code class="code docutils literal notranslate"><span class="pre">RecordToFile</span></code> to dump measurement records into a file <cite>conv2d.json</cite>.
The measurement records can be used to query the history best, resume the search,
and do more analyses later.</p></li>
<li><p>see <a class="reference internal" href="../../reference/api/python/auto_scheduler.html#tvm.auto_scheduler.TuningOptions" title="tvm.auto_scheduler.TuningOptions"><code class="xref any py py-class docutils literal notranslate"><span class="pre">auto_scheduler.TuningOptions</span></code></a>,
<a class="reference internal" href="../../reference/api/python/auto_scheduler.html#tvm.auto_scheduler.LocalRPCMeasureContext" title="tvm.auto_scheduler.LocalRPCMeasureContext"><code class="xref any py py-class docutils literal notranslate"><span class="pre">auto_scheduler.LocalRPCMeasureContext</span></code></a> for more parameters.</p></li>
</ul>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">log_file</span> <span class="o">=</span> <span class="s2">&quot;conv2d.json&quot;</span>
<span class="n">measure_ctx</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">LocalRPCMeasureContext</span><span class="p">(</span><span class="n">min_repeat_ms</span><span class="o">=</span><span class="mi">300</span><span class="p">)</span>
<span class="n">tune_option</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">TuningOptions</span><span class="p">(</span>
    <span class="n">num_measure_trials</span><span class="o">=</span><span class="mi">10</span><span class="p">,</span>  <span class="c1"># change this to 1000 to achieve the best performance</span>
    <span class="n">runner</span><span class="o">=</span><span class="n">measure_ctx</span><span class="o">.</span><span class="n">runner</span><span class="p">,</span>
    <span class="n">measure_callbacks</span><span class="o">=</span><span class="p">[</span><span class="n">auto_scheduler</span><span class="o">.</span><span class="n">RecordToFile</span><span class="p">(</span><span class="n">log_file</span><span class="p">)],</span>
    <span class="n">verbose</span><span class="o">=</span><span class="mi">2</span><span class="p">,</span>
<span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Get devices for measurement successfully!
</pre></div>
</div>
</div>
<div class="section" id="run-the-search">
<h2>Run the search<a class="headerlink" href="#run-the-search" title="永久链接至标题">¶</a></h2>
<p>Now we get all inputs ready. Pretty simple, isn’t it?
We can kick off the search and let the auto-scheduler do its magic.
After some measurement trials, we can load the best schedule from the log
file and apply it.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># Run auto-tuning (search)</span>
<span class="n">task</span><span class="o">.</span><span class="n">tune</span><span class="p">(</span><span class="n">tune_option</span><span class="p">)</span>
<span class="c1"># Apply the best schedule</span>
<span class="n">sch</span><span class="p">,</span> <span class="n">args</span> <span class="o">=</span> <span class="n">task</span><span class="o">.</span><span class="n">apply_best</span><span class="p">(</span><span class="n">log_file</span><span class="p">)</span>

<span class="c1"># Kill the measurement process</span>
<span class="k">del</span> <span class="n">measure_ctx</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>
</pre></div>
</div>
<p>We can lower the schedule to see the IR after auto-scheduling.
The auto-scheduler correctly performs optimizations including multi-level tiling,
cooperative fetching, unrolling and operator fusion.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Lowered TIR:&quot;</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">tvm</span><span class="o">.</span><span class="n">lower</span><span class="p">(</span><span class="n">sch</span><span class="p">,</span> <span class="n">args</span><span class="p">,</span> <span class="n">simple_mode</span><span class="o">=</span><span class="kc">True</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Lowered TIR:
primfn(data_1: handle, kernel_1: handle, bias_1: handle, compute_1: handle) -&gt; ()
  attr = {&quot;from_legacy_te_schedule&quot;: True, &quot;global_symbol&quot;: &quot;main&quot;, &quot;tir.noalias&quot;: True}
  buffers = {compute: Buffer(compute_2: Pointer(float32), float32, [1, 512, 7, 7], []),
             data: Buffer(data_2: Pointer(float32), float32, [1, 512, 7, 7], []),
             bias: Buffer(bias_2: Pointer(float32), float32, [1, 512, 1, 1], []),
             kernel: Buffer(kernel_2: Pointer(float32), float32, [512, 512, 3, 3], [])}
  buffer_map = {data_1: data, kernel_1: kernel, bias_1: bias, compute_1: compute} {
  attr [IterVar(blockIdx.x: int32, (nullptr), &quot;ThreadIndex&quot;, &quot;blockIdx.x&quot;)] &quot;thread_extent&quot; = 32;
  allocate(compute_3: Pointer(local float32), float32, [14]), storage_scope = local;
  allocate(pad_temp.shared: Pointer(shared float32), float32, [648]), storage_scope = shared;
  allocate(kernel.shared: Pointer(shared float32), float32, [1152]), storage_scope = shared;
  attr [IterVar(threadIdx.x: int32, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56 {
    compute_3[0] = 0f32
    compute_3[1] = 0f32
    compute_3[2] = 0f32
    compute_3[3] = 0f32
    compute_3[4] = 0f32
    compute_3[5] = 0f32
    compute_3[6] = 0f32
    compute_3[7] = 0f32
    compute_3[8] = 0f32
    compute_3[9] = 0f32
    compute_3[10] = 0f32
    compute_3[11] = 0f32
    compute_3[12] = 0f32
    compute_3[13] = 0f32
    for (rc.outer.outer: int32, 0, 64) {
      attr [IterVar(threadIdx.x_1: int32, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56 {
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[(threadIdx.x_1*12)] = @tir.if_then_else(((((9 &lt;= floormod((threadIdx.x_1*12), 81)) &amp;&amp; (floormod((threadIdx.x_1*12), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod((threadIdx.x_1*12), 9))) &amp;&amp; (floormod((threadIdx.x_1*12), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv((threadIdx.x_1*12), 81)*49)) + (floordiv(floormod((threadIdx.x_1*12), 81), 9)*7)) + floormod((threadIdx.x_1*12), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 1)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 1), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 1), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 1), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 1), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 1), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 1), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 1), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 2)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 2), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 2), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 2), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 2), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 2), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 2), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 2), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 3)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 3), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 3), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 3), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 3), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 3), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 3), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 3), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 4)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 4), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 4), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 4), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 4), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 4), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 4), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 4), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 5)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 5), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 5), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 5), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 5), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 5), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 5), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 5), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 6)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 6), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 6), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 6), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 6), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 6), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 6), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 6), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 7)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 7), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 7), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 7), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 7), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 7), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 7), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 7), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 8)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 8), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 8), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 8), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 8), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 8), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 8), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 8), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 9)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 9), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 9), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod((threadIdx.x_1*12), 9))) &amp;&amp; (floormod((threadIdx.x_1*12), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 9), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 9), 81), 9)*7)) + floormod((threadIdx.x_1*12), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 10)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 10), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 10), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 1), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 1), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 10), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 10), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 1), 9)) - 8)], 0f32, dtype=float32)
        }
        if @tir.likely((threadIdx.x_1 &lt; 54), dtype=bool) {
          pad_temp.shared[((threadIdx.x_1*12) + 11)] = @tir.if_then_else(((((9 &lt;= floormod(((threadIdx.x_1*12) + 11), 81)) &amp;&amp; (floormod(((threadIdx.x_1*12) + 11), 81) &lt; 72)) &amp;&amp; (1 &lt;= floormod(((threadIdx.x_1*12) + 2), 9))) &amp;&amp; (floormod(((threadIdx.x_1*12) + 2), 9) &lt; 8)), (float32*)data_2[(((((rc.outer.outer*392) + (floordiv(((threadIdx.x_1*12) + 11), 81)*49)) + (floordiv(floormod(((threadIdx.x_1*12) + 11), 81), 9)*7)) + floormod(((threadIdx.x_1*12) + 2), 9)) - 8)], 0f32, dtype=float32)
        }
      }
      attr [IterVar(threadIdx.x_2: int32, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[threadIdx.x_2] = (float32*)kernel_2[(((blockIdx.x*73728) + (rc.outer.outer*72)) + threadIdx.x_2)]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 56)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 56), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 56), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 112)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 112), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 40), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 168)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 168), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 24), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 224)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 224), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 8), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 280)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 280), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 64), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 336)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 336), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 48), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 392)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 392), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 32), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 448)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 448), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 16), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 504)] = (float32*)kernel_2[((((blockIdx.x*73728) + (rc.outer.outer*72)) + threadIdx.x_2) + 32256)]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 560)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 560), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 56), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 616)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 616), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 40), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 672)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 672), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 24), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 728)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 728), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 8), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 784)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 784), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 64), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 840)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 840), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 48), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 896)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 896), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 32), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 952)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 952), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 16), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 1008)] = (float32*)kernel_2[((((blockIdx.x*73728) + (rc.outer.outer*72)) + threadIdx.x_2) + 64512)]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      kernel.shared[(threadIdx.x_2 + 1064)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 1064), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 56), 72))]
      attr [IterVar(threadIdx.x_2, (nullptr), &quot;ThreadIndex&quot;, &quot;threadIdx.x&quot;)] &quot;thread_extent&quot; = 56;
      if @tir.likely((threadIdx.x_2 &lt; 32), dtype=bool) {
        kernel.shared[(threadIdx.x_2 + 1120)] = (float32*)kernel_2[((((blockIdx.x*73728) + (floordiv((threadIdx.x_2 + 1120), 72)*4608)) + (rc.outer.outer*72)) + floormod((threadIdx.x_2 + 40), 72))]
      }
      for (rc.outer.inner: int32, 0, 8) {
        for (rx.outer.inner: int32, 0, 3) {
          compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[(((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7))]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 9)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[(((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner)]))
          compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[(((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7))]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 9)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 72)]))
          compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 9)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 63)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 3)]))
          compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 9)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 63)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 75)]))
          compute_3[0] = ((float32*)compute_3[0] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[1] = ((float32*)compute_3[1] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[2] = ((float32*)compute_3[2] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[3] = ((float32*)compute_3[3] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[4] = ((float32*)compute_3[4] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[5] = ((float32*)compute_3[5] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 63)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[6] = ((float32*)compute_3[6] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 72)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 6)]))
          compute_3[7] = ((float32*)compute_3[7] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 18)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[8] = ((float32*)compute_3[8] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 27)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[9] = ((float32*)compute_3[9] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 36)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[10] = ((float32*)compute_3[10] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 45)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[11] = ((float32*)compute_3[11] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 54)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[12] = ((float32*)compute_3[12] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 63)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
          compute_3[13] = ((float32*)compute_3[13] + ((float32*)pad_temp.shared[((((rc.outer.inner*81) + rx.outer.inner) + floormod(threadIdx.x, 7)) + 72)]*(float32*)kernel.shared[((((floordiv(threadIdx.x, 7)*144) + (rc.outer.inner*9)) + rx.outer.inner) + 78)]))
        }
      }
    }
    for (i1.inner: int32, 0, 2) {
      for (i2.inner: int32, 0, 7) {
        compute_2[(((((blockIdx.x*784) + (floordiv(threadIdx.x, 7)*98)) + (i1.inner*49)) + (i2.inner*7)) + floormod(threadIdx.x, 7))] = max(((float32*)compute_3[((i1.inner*7) + i2.inner)] + (float32*)bias_2[(((blockIdx.x*16) + (floordiv(threadIdx.x, 7)*2)) + i1.inner)]), 0f32)
      }
    }
  }
}
</pre></div>
</div>
</div>
<div class="section" id="check-correctness-and-evaluate-performance">
<h2>Check correctness and evaluate performance<a class="headerlink" href="#check-correctness-and-evaluate-performance" title="永久链接至标题">¶</a></h2>
<p>We build the binary and check its correctness and performance.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">func</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">sch</span><span class="p">,</span> <span class="n">args</span><span class="p">,</span> <span class="n">target</span><span class="p">)</span>

<span class="c1"># Check correctness</span>
<span class="n">data_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)</span>
<span class="n">weight_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)</span>
<span class="n">bias_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="mi">1</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)</span>
<span class="n">conv_np</span> <span class="o">=</span> <span class="n">conv2d_nchw_python</span><span class="p">(</span><span class="n">data_np</span><span class="p">,</span> <span class="n">weight_np</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span><span class="p">)</span>
<span class="n">out_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">maximum</span><span class="p">(</span><span class="n">conv_np</span> <span class="o">+</span> <span class="n">bias_np</span><span class="p">,</span> <span class="mf">0.0</span><span class="p">)</span>

<span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">cuda</span><span class="p">()</span>
<span class="n">data_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">data_np</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">weight_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">weight_np</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">bias_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">bias_np</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">out_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">empty</span><span class="p">(</span><span class="n">out_np</span><span class="o">.</span><span class="n">shape</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">func</span><span class="p">(</span><span class="n">data_tvm</span><span class="p">,</span> <span class="n">weight_tvm</span><span class="p">,</span> <span class="n">bias_tvm</span><span class="p">,</span> <span class="n">out_tvm</span><span class="p">)</span>

<span class="c1"># Check results</span>
<span class="n">np</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">out_np</span><span class="p">,</span> <span class="n">out_tvm</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">rtol</span><span class="o">=</span><span class="mf">1e-3</span><span class="p">)</span>

<span class="c1"># Evaluate execution time</span>
<span class="n">evaluator</span> <span class="o">=</span> <span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span><span class="p">(</span><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">min_repeat_ms</span><span class="o">=</span><span class="mi">500</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span>
    <span class="s2">&quot;Execution time of this operator: </span><span class="si">%.3f</span><span class="s2"> ms&quot;</span>
    <span class="o">%</span> <span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">median</span><span class="p">(</span><span class="n">evaluator</span><span class="p">(</span><span class="n">data_tvm</span><span class="p">,</span> <span class="n">weight_tvm</span><span class="p">,</span> <span class="n">bias_tvm</span><span class="p">,</span> <span class="n">out_tvm</span><span class="p">)</span><span class="o">.</span><span class="n">results</span><span class="p">)</span> <span class="o">*</span> <span class="mi">1000</span><span class="p">)</span>
<span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Execution time of this operator: 0.164 ms
</pre></div>
</div>
</div>
<div class="section" id="using-the-record-file">
<h2>Using the record file<a class="headerlink" href="#using-the-record-file" title="永久链接至标题">¶</a></h2>
<p>During the search, all measurement records are dumped into the record
file “conv2d.json”. The measurement records can be used to re-apply search results,
resume the search, and perform other analyses.</p>
<p>Here is an example where we load the best schedule from a file,
print the equivalent python schedule API and CUDA source code.
They can be used for debugging and learning the behavior of the auto-scheduler.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Equivalent python schedule:&quot;</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">task</span><span class="o">.</span><span class="n">print_best</span><span class="p">(</span><span class="n">log_file</span><span class="p">,</span> <span class="n">print_mode</span><span class="o">=</span><span class="s2">&quot;schedule&quot;</span><span class="p">))</span>

<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;CUDA source code:&quot;</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">task</span><span class="o">.</span><span class="n">print_best</span><span class="p">(</span><span class="n">log_file</span><span class="p">,</span> <span class="n">print_mode</span><span class="o">=</span><span class="s2">&quot;cuda&quot;</span><span class="p">))</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Equivalent python schedule:
pad_temp_i0, pad_temp_i1, pad_temp_i2, pad_temp_i3 = tuple(pad_temp.op.axis) + tuple(pad_temp.op.reduce_axis)
compute_nn, compute_ff, compute_yy, compute_xx, compute_rc, compute_ry, compute_rx = tuple(compute.op.axis) + tuple(compute.op.reduce_axis)
T_add_ax0, T_add_ax1, T_add_ax2, T_add_ax3 = tuple(T_add.op.axis) + tuple(T_add.op.reduce_axis)
compute_i0, compute_i1, compute_i2, compute_i3 = tuple(compute.op.axis) + tuple(compute.op.reduce_axis)
s[T_add].compute_inline()
compute_nn_o_i, compute_nn_i = s[compute].split(compute_nn, factor=1)
compute_nn_o_o_i, compute_nn_o_i = s[compute].split(compute_nn_o_i, factor=1)
compute_nn_o_o_o_i, compute_nn_o_o_i = s[compute].split(compute_nn_o_o_i, factor=1)
compute_nn_o_o_o_o, compute_nn_o_o_o_i = s[compute].split(compute_nn_o_o_o_i, factor=1)
compute_ff_o_i, compute_ff_i = s[compute].split(compute_ff, factor=2)
compute_ff_o_o_i, compute_ff_o_i = s[compute].split(compute_ff_o_i, factor=1)
compute_ff_o_o_o_i, compute_ff_o_o_i = s[compute].split(compute_ff_o_o_i, factor=8)
compute_ff_o_o_o_o, compute_ff_o_o_o_i = s[compute].split(compute_ff_o_o_o_i, factor=1)
compute_yy_o_i, compute_yy_i = s[compute].split(compute_yy, factor=7)
compute_yy_o_o_i, compute_yy_o_i = s[compute].split(compute_yy_o_i, factor=1)
compute_yy_o_o_o_i, compute_yy_o_o_i = s[compute].split(compute_yy_o_o_i, factor=1)
compute_yy_o_o_o_o, compute_yy_o_o_o_i = s[compute].split(compute_yy_o_o_o_i, factor=1)
compute_xx_o_i, compute_xx_i = s[compute].split(compute_xx, factor=1)
compute_xx_o_o_i, compute_xx_o_i = s[compute].split(compute_xx_o_i, factor=1)
compute_xx_o_o_o_i, compute_xx_o_o_i = s[compute].split(compute_xx_o_o_i, factor=7)
compute_xx_o_o_o_o, compute_xx_o_o_o_i = s[compute].split(compute_xx_o_o_o_i, factor=1)
compute_rc_o_i, compute_rc_i = s[compute].split(compute_rc, factor=1)
compute_rc_o_o, compute_rc_o_i = s[compute].split(compute_rc_o_i, factor=8)
compute_ry_o_i, compute_ry_i = s[compute].split(compute_ry, factor=3)
compute_ry_o_o, compute_ry_o_i = s[compute].split(compute_ry_o_i, factor=1)
compute_rx_o_i, compute_rx_i = s[compute].split(compute_rx, factor=1)
compute_rx_o_o, compute_rx_o_i = s[compute].split(compute_rx_o_i, factor=3)
s[compute].reorder(compute_nn_o_o_o_o, compute_ff_o_o_o_o, compute_yy_o_o_o_o, compute_xx_o_o_o_o, compute_nn_o_o_o_i, compute_ff_o_o_o_i, compute_yy_o_o_o_i, compute_xx_o_o_o_i, compute_nn_o_o_i, compute_ff_o_o_i, compute_yy_o_o_i, compute_xx_o_o_i, compute_rc_o_o, compute_ry_o_o, compute_rx_o_o, compute_rc_o_i, compute_ry_o_i, compute_rx_o_i, compute_nn_o_i, compute_ff_o_i, compute_yy_o_i, compute_xx_o_i, compute_rc_i, compute_ry_i, compute_rx_i, compute_nn_i, compute_ff_i, compute_yy_i, compute_xx_i)
compute_i0_o_i, compute_i0_i = s[compute].split(compute_i0, factor=1)
compute_i0_o_o_i, compute_i0_o_i = s[compute].split(compute_i0_o_i, factor=1)
compute_i0_o_o_o, compute_i0_o_o_i = s[compute].split(compute_i0_o_o_i, factor=1)
compute_i1_o_i, compute_i1_i = s[compute].split(compute_i1, factor=2)
compute_i1_o_o_i, compute_i1_o_i = s[compute].split(compute_i1_o_i, factor=8)
compute_i1_o_o_o, compute_i1_o_o_i = s[compute].split(compute_i1_o_o_i, factor=1)
compute_i2_o_i, compute_i2_i = s[compute].split(compute_i2, factor=7)
compute_i2_o_o_i, compute_i2_o_i = s[compute].split(compute_i2_o_i, factor=1)
compute_i2_o_o_o, compute_i2_o_o_i = s[compute].split(compute_i2_o_o_i, factor=1)
compute_i3_o_i, compute_i3_i = s[compute].split(compute_i3, factor=1)
compute_i3_o_o_i, compute_i3_o_i = s[compute].split(compute_i3_o_i, factor=7)
compute_i3_o_o_o, compute_i3_o_o_i = s[compute].split(compute_i3_o_o_i, factor=1)
s[compute].reorder(compute_i0_o_o_o, compute_i1_o_o_o, compute_i2_o_o_o, compute_i3_o_o_o, compute_i0_o_o_i, compute_i1_o_o_i, compute_i2_o_o_i, compute_i3_o_o_i, compute_i0_o_i, compute_i1_o_i, compute_i2_o_i, compute_i3_o_i, compute_i0_i, compute_i1_i, compute_i2_i, compute_i3_i)
s[compute].compute_at(s[compute], compute_i3_o_i)
kernel_shared = s.cache_read(kernel, &quot;shared&quot;, [compute])
kernel_shared_ax0, kernel_shared_ax1, kernel_shared_ax2, kernel_shared_ax3 = tuple(kernel_shared.op.axis)
s[kernel_shared].compute_at(s[compute], compute_rx_o_o)
pad_temp_shared = s.cache_read(pad_temp, &quot;shared&quot;, [compute])
pad_temp_shared_ax0, pad_temp_shared_ax1, pad_temp_shared_ax2, pad_temp_shared_ax3 = tuple(pad_temp_shared.op.axis)
s[pad_temp_shared].compute_at(s[compute], compute_rx_o_o)
s[pad_temp].compute_inline()
compute_i0_o_o_o_i1_o_o_o_fused_i2_o_o_o_fused_i3_o_o_o_fused = s[compute].fuse(compute_i0_o_o_o, compute_i1_o_o_o, compute_i2_o_o_o, compute_i3_o_o_o)
s[compute].bind(compute_i0_o_o_o_i1_o_o_o_fused_i2_o_o_o_fused_i3_o_o_o_fused, te.thread_axis(&quot;blockIdx.x&quot;))
compute_i0_o_o_i_i1_o_o_i_fused_i2_o_o_i_fused_i3_o_o_i_fused = s[compute].fuse(compute_i0_o_o_i, compute_i1_o_o_i, compute_i2_o_o_i, compute_i3_o_o_i)
s[compute].bind(compute_i0_o_o_i_i1_o_o_i_fused_i2_o_o_i_fused_i3_o_o_i_fused, te.thread_axis(&quot;vthread&quot;))
compute_i0_o_i_i1_o_i_fused_i2_o_i_fused_i3_o_i_fused = s[compute].fuse(compute_i0_o_i, compute_i1_o_i, compute_i2_o_i, compute_i3_o_i)
s[compute].bind(compute_i0_o_i_i1_o_i_fused_i2_o_i_fused_i3_o_i_fused, te.thread_axis(&quot;threadIdx.x&quot;))
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused = s[kernel_shared].fuse(kernel_shared_ax0, kernel_shared_ax1, kernel_shared_ax2, kernel_shared_ax3)
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i = s[kernel_shared].split(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused, factor=1)
s[kernel_shared].vectorize(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i)
kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_o, kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i = s[kernel_shared].split(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, factor=56)
s[kernel_shared].bind(kernel_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i, te.thread_axis(&quot;threadIdx.x&quot;))
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused = s[pad_temp_shared].fuse(pad_temp_shared_ax0, pad_temp_shared_ax1, pad_temp_shared_ax2, pad_temp_shared_ax3)
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i = s[pad_temp_shared].split(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused, factor=12)
s[pad_temp_shared].vectorize(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_i)
pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_o, pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i = s[pad_temp_shared].split(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o, factor=56)
s[pad_temp_shared].bind(pad_temp_shared_ax0_ax1_fused_ax2_fused_ax3_fused_o_i, te.thread_axis(&quot;threadIdx.x&quot;))
s[compute].pragma(compute_nn_o_o_o_o, &quot;auto_unroll_max_step&quot;, 64)
s[compute].pragma(compute_nn_o_o_o_o, &quot;unroll_explicit&quot;, True)

CUDA source code:

#ifdef _WIN32
  using uint = unsigned int;
  using uchar = unsigned char;
  using ushort = unsigned short;
  using int64_t = long long;
  using uint64_t = unsigned long long;
#else
  #define uint unsigned int
  #define uchar unsigned char
  #define ushort unsigned short
  #define int64_t long long
  #define uint64_t unsigned long long
#endif
extern &quot;C&quot; __global__ void __launch_bounds__(56) default_function_kernel0(float* __restrict__ data, float* __restrict__ kernel, float* __restrict__ compute, float* __restrict__ bias) {
  float compute1[14];
  __shared__ float pad_temp_shared[648];
  __shared__ float kernel_shared[1152];
  compute1[(0)] = 0.000000e+00f;
  compute1[(1)] = 0.000000e+00f;
  compute1[(2)] = 0.000000e+00f;
  compute1[(3)] = 0.000000e+00f;
  compute1[(4)] = 0.000000e+00f;
  compute1[(5)] = 0.000000e+00f;
  compute1[(6)] = 0.000000e+00f;
  compute1[(7)] = 0.000000e+00f;
  compute1[(8)] = 0.000000e+00f;
  compute1[(9)] = 0.000000e+00f;
  compute1[(10)] = 0.000000e+00f;
  compute1[(11)] = 0.000000e+00f;
  compute1[(12)] = 0.000000e+00f;
  compute1[(13)] = 0.000000e+00f;
  for (int rc_outer_outer = 0; rc_outer_outer &lt; 64; ++rc_outer_outer) {
    __syncthreads();
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[((((int)threadIdx.x) * 12))] = (((((9 &lt;= ((((int)threadIdx.x) * 12) % 81)) &amp;&amp; (((((int)threadIdx.x) * 12) % 81) &lt; 72)) &amp;&amp; (1 &lt;= ((((int)threadIdx.x) * 12) % 9))) &amp;&amp; (((((int)threadIdx.x) * 12) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + (((((int)threadIdx.x) * 12) / 81) * 49)) + ((((((int)threadIdx.x) * 12) % 81) / 9) * 7)) + ((((int)threadIdx.x) * 12) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 1))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 1) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 1) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 1) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 1) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 1) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 1) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 1) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 2))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 2) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 2) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 2) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 2) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 2) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 2) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 2) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 3))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 3) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 3) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 3) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 3) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 3) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 3) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 3) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 4))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 4) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 4) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 4) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 4) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 4) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 4) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 4) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 5))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 5) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 5) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 5) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 5) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 5) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 5) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 5) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 6))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 6) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 6) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 6) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 6) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 6) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 6) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 6) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 7))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 7) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 7) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 7) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 7) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 7) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 7) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 7) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 8))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 8) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 8) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 8) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 8) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 8) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 8) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 8) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 9))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 9) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 9) % 81) &lt; 72)) &amp;&amp; (1 &lt;= ((((int)threadIdx.x) * 12) % 9))) &amp;&amp; (((((int)threadIdx.x) * 12) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 9) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 9) % 81) / 9) * 7)) + ((((int)threadIdx.x) * 12) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 10))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 10) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 10) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 1) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 1) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 10) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 10) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 1) % 9)) - 8))] : 0.000000e+00f);
    }
    if (((int)threadIdx.x) &lt; 54) {
      pad_temp_shared[(((((int)threadIdx.x) * 12) + 11))] = (((((9 &lt;= (((((int)threadIdx.x) * 12) + 11) % 81)) &amp;&amp; ((((((int)threadIdx.x) * 12) + 11) % 81) &lt; 72)) &amp;&amp; (1 &lt;= (((((int)threadIdx.x) * 12) + 2) % 9))) &amp;&amp; ((((((int)threadIdx.x) * 12) + 2) % 9) &lt; 8)) ? data[((((((rc_outer_outer * 392) + ((((((int)threadIdx.x) * 12) + 11) / 81) * 49)) + (((((((int)threadIdx.x) * 12) + 11) % 81) / 9) * 7)) + (((((int)threadIdx.x) * 12) + 2) % 9)) - 8))] : 0.000000e+00f);
    }
    kernel_shared[(((int)threadIdx.x))] = kernel[((((((int)blockIdx.x) * 73728) + (rc_outer_outer * 72)) + ((int)threadIdx.x)))];
    kernel_shared[((((int)threadIdx.x) + 56))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 56) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 56) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 112))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 112) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 40) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 168))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 168) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 24) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 224))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 224) / 72) * 4608)) + (rc_outer_outer * 72)) + (((int)threadIdx.x) + 8)))];
    kernel_shared[((((int)threadIdx.x) + 280))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 280) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 64) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 336))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 336) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 48) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 392))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 392) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 32) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 448))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 448) / 72) * 4608)) + (rc_outer_outer * 72)) + (((int)threadIdx.x) + 16)))];
    kernel_shared[((((int)threadIdx.x) + 504))] = kernel[(((((((int)blockIdx.x) * 73728) + (rc_outer_outer * 72)) + ((int)threadIdx.x)) + 32256))];
    kernel_shared[((((int)threadIdx.x) + 560))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 560) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 56) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 616))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 616) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 40) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 672))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 672) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 24) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 728))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 728) / 72) * 4608)) + (rc_outer_outer * 72)) + (((int)threadIdx.x) + 8)))];
    kernel_shared[((((int)threadIdx.x) + 784))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 784) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 64) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 840))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 840) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 48) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 896))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 896) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 32) % 72)))];
    kernel_shared[((((int)threadIdx.x) + 952))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 952) / 72) * 4608)) + (rc_outer_outer * 72)) + (((int)threadIdx.x) + 16)))];
    kernel_shared[((((int)threadIdx.x) + 1008))] = kernel[(((((((int)blockIdx.x) * 73728) + (rc_outer_outer * 72)) + ((int)threadIdx.x)) + 64512))];
    kernel_shared[((((int)threadIdx.x) + 1064))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 1064) / 72) * 4608)) + (rc_outer_outer * 72)) + ((((int)threadIdx.x) + 56) % 72)))];
    if (((int)threadIdx.x) &lt; 32) {
      kernel_shared[((((int)threadIdx.x) + 1120))] = kernel[(((((((int)blockIdx.x) * 73728) + (((((int)threadIdx.x) + 1120) / 72) * 4608)) + (rc_outer_outer * 72)) + (((int)threadIdx.x) + 40)))];
    }
    __syncthreads();
    for (int rc_outer_inner = 0; rc_outer_inner &lt; 8; ++rc_outer_inner) {
      for (int rx_outer_inner = 0; rx_outer_inner &lt; 3; ++rx_outer_inner) {
        compute1[(0)] = (compute1[(0)] + (pad_temp_shared[((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(1)] = (compute1[(1)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 9))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(2)] = (compute1[(2)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(3)] = (compute1[(3)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(4)] = (compute1[(4)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(5)] = (compute1[(5)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(6)] = (compute1[(6)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[(((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner))]));
        compute1[(7)] = (compute1[(7)] + (pad_temp_shared[((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(8)] = (compute1[(8)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 9))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(9)] = (compute1[(9)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(10)] = (compute1[(10)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(11)] = (compute1[(11)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(12)] = (compute1[(12)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(13)] = (compute1[(13)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 72))]));
        compute1[(0)] = (compute1[(0)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 9))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(1)] = (compute1[(1)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(2)] = (compute1[(2)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(3)] = (compute1[(3)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(4)] = (compute1[(4)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(5)] = (compute1[(5)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(6)] = (compute1[(6)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 63))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 3))]));
        compute1[(7)] = (compute1[(7)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 9))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(8)] = (compute1[(8)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(9)] = (compute1[(9)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(10)] = (compute1[(10)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(11)] = (compute1[(11)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(12)] = (compute1[(12)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(13)] = (compute1[(13)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 63))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 75))]));
        compute1[(0)] = (compute1[(0)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(1)] = (compute1[(1)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(2)] = (compute1[(2)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(3)] = (compute1[(3)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(4)] = (compute1[(4)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(5)] = (compute1[(5)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 63))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(6)] = (compute1[(6)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 72))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 6))]));
        compute1[(7)] = (compute1[(7)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 18))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(8)] = (compute1[(8)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 27))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(9)] = (compute1[(9)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 36))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(10)] = (compute1[(10)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 45))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(11)] = (compute1[(11)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 54))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(12)] = (compute1[(12)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 63))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
        compute1[(13)] = (compute1[(13)] + (pad_temp_shared[(((((rc_outer_inner * 81) + rx_outer_inner) + (((int)threadIdx.x) % 7)) + 72))] * kernel_shared[((((((((int)threadIdx.x) / 7) * 144) + (rc_outer_inner * 9)) + rx_outer_inner) + 78))]));
      }
    }
  }
  for (int i1_inner = 0; i1_inner &lt; 2; ++i1_inner) {
    for (int i2_inner = 0; i2_inner &lt; 7; ++i2_inner) {
      compute[((((((((int)blockIdx.x) * 784) + ((((int)threadIdx.x) / 7) * 98)) + (i1_inner * 49)) + (i2_inner * 7)) + (((int)threadIdx.x) % 7)))] = max((compute1[(((i1_inner * 7) + i2_inner))] + bias[((((((int)blockIdx.x) * 16) + ((((int)threadIdx.x) / 7) * 2)) + i1_inner))]), 0.000000e+00f);
    }
  }
}
</pre></div>
</div>
<p>A more complicated example is to resume the search.
In this case, we need to create the search policy and cost model by ourselves
and resume the status of search policy and cost model with the log file.
In the example below we resume the status and do more 5 trials.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="k">def</span> <span class="nf">resume_search</span><span class="p">(</span><span class="n">task</span><span class="p">,</span> <span class="n">log_file</span><span class="p">):</span>
    <span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Resume search:&quot;</span><span class="p">)</span>
    <span class="n">cost_model</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">XGBModel</span><span class="p">()</span>
    <span class="n">cost_model</span><span class="o">.</span><span class="n">update_from_file</span><span class="p">(</span><span class="n">log_file</span><span class="p">)</span>
    <span class="n">search_policy</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">SketchPolicy</span><span class="p">(</span>
        <span class="n">task</span><span class="p">,</span> <span class="n">cost_model</span><span class="p">,</span> <span class="n">init_search_callbacks</span><span class="o">=</span><span class="p">[</span><span class="n">auto_scheduler</span><span class="o">.</span><span class="n">PreloadMeasuredStates</span><span class="p">(</span><span class="n">log_file</span><span class="p">)]</span>
    <span class="p">)</span>
    <span class="n">measure_ctx</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">LocalRPCMeasureContext</span><span class="p">(</span><span class="n">min_repeat_ms</span><span class="o">=</span><span class="mi">300</span><span class="p">)</span>
    <span class="n">tune_option</span> <span class="o">=</span> <span class="n">auto_scheduler</span><span class="o">.</span><span class="n">TuningOptions</span><span class="p">(</span>
        <span class="n">num_measure_trials</span><span class="o">=</span><span class="mi">5</span><span class="p">,</span>
        <span class="n">runner</span><span class="o">=</span><span class="n">measure_ctx</span><span class="o">.</span><span class="n">runner</span><span class="p">,</span>
        <span class="n">measure_callbacks</span><span class="o">=</span><span class="p">[</span><span class="n">auto_scheduler</span><span class="o">.</span><span class="n">RecordToFile</span><span class="p">(</span><span class="n">log_file</span><span class="p">)],</span>
    <span class="p">)</span>
    <span class="n">task</span><span class="o">.</span><span class="n">tune</span><span class="p">(</span><span class="n">tune_option</span><span class="p">,</span> <span class="n">search_policy</span><span class="o">=</span><span class="n">search_policy</span><span class="p">)</span>

    <span class="c1"># Kill the measurement process</span>
    <span class="k">del</span> <span class="n">measure_ctx</span>


<span class="n">resume_search</span><span class="p">(</span><span class="n">task</span><span class="p">,</span> <span class="n">log_file</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Resume search:
/usr/local/lib/python3.6/dist-packages/xgboost/training.py:17: UserWarning: Old style callback is deprecated.  See: https://xgboost.readthedocs.io/en/latest/python/callbacks.html
  warnings.warn(f&#39;Old style callback is deprecated.  See: {link}&#39;, UserWarning)
Get devices for measurement successfully!
</pre></div>
</div>
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 1 minutes  52.623 seconds)</p>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-how-to-tune-with-autoscheduler-tune-conv2d-layer-cuda-py">
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/e3e540f3b477c0c52d8eb73e674e8ffd/tune_conv2d_layer_cuda.py"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Python</span> <span class="pre">source</span> <span class="pre">code:</span> <span class="pre">tune_conv2d_layer_cuda.py</span></code></a></p>
</div>
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/5f1f7bd7d90710fd404f7bcdc4965622/tune_conv2d_layer_cuda.ipynb"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Jupyter</span> <span class="pre">notebook:</span> <span class="pre">tune_conv2d_layer_cuda.ipynb</span></code></a></p>
</div>
</div>
<p class="sphx-glr-signature"><a class="reference external" href="https://sphinx-gallery.github.io">Gallery generated by Sphinx-Gallery</a></p>
</div>
</div>


           </div>
           
          </div>
          

<footer>

    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="tune_network_x86.html" class="btn btn-neutral float-right" title="Auto-scheduling a Neural Network for x86 CPU" accesskey="n" rel="next">下一个 <span class="fa fa-arrow-circle-right"></span></a>
      
      
        <a href="index.html" class="btn btn-neutral float-left" title="Use AutoScheduler for Template-Free Scheduling" accesskey="p" rel="prev"><span class="fa fa-arrow-circle-left"></span> 上一个</a>
      
    </div>

<div id="button" class="backtop"><img src="../../_static//img/right.svg" alt="backtop"/> </div>
<section class="footerSec">
    <div class="footerHeader">
      <ul class="d-flex align-md-items-center justify-content-between flex-column flex-md-row">
        <li class="copywrite d-flex align-items-center">
          <h5 id="copy-right-info">© 2020 Apache Software Foundation | All right reserved</h5>
        </li>
      </ul>

    </div>

    <ul>
      <li class="footernote">Copyright © 2020 The Apache Software Foundation. Apache TVM, Apache, the Apache feather, and the Apache TVM project logo are either trademarks or registered trademarks of the Apache Software Foundation.</li>
    </ul>

</section>
</footer>
        </div>
      </div>

    </section>

  </div>
  

    <script src="https://cdnjs.cloudflare.com/ajax/libs/popper.js/1.12.9/umd/popper.min.js" integrity="sha384-ApNbgh9B+Y1QKtv3Rn7W3mgPxhU9K/ScQsAP7hUibX39j7fakFPskvXusvfa0b4Q" crossorigin="anonymous"></script>
    <script src="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/js/bootstrap.min.js" integrity="sha384-JZR6Spejh4U02d8jOt6vLEHfe/JQGiRRSQQxSfFWpi1MquVdAyjUar5+76PVCmYl" crossorigin="anonymous"></script>

  </body>
  <script type="text/javascript">
      jQuery(function () {
          SphinxRtdTheme.Navigation.enable(true);
      });
  </script>

  
  
    
    <!-- Theme Analytics -->
    <script>
    (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){
      (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o),
      m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m)
    })(window,document,'script','https://www.google-analytics.com/analytics.js','ga');

    ga('create', 'UA-75982049-2', 'auto');
    ga('send', 'pageview');
    </script>

    
   

</body>
</html>